#include "matrix3D.cuh"
#include <cstdio>
#include "timer.cuh"
#include "cuda_runtime.h"

__global__ void transform_kernel(const short * __restrict__ source,short* __restrict__ dist, const int* classes, int classesCount)
{
     unsigned int thread = blockDim.x * blockIdx.x + threadIdx.x;
     short val = 0;
     int sourceInRange = (source[thread] + 32768) >> 8;
     for(int k = 0 ; k < classesCount-1; k++)
     {
         if(sourceInRange >= classes[k] && sourceInRange <= classes[k+1])
         {
             val = k << 4;
             k = classesCount;
         }
     }
     dist[thread] = val;
}



void cuda_classify(const short *source,short *dist, int sourceSize, const int* classes, int classesCount)
{
    short* d_source;
    short* d_dist;
    int* d_classes;

    cudaError_t cudaStatus;

    cudaStatus = cudaMalloc((void**)&d_source, sourceSize * sizeof(short));
    if (cudaStatus != cudaSuccess)
    {
        printf("cudaMalloc failed for model memory request!\n");
        return;
    }

    cudaStatus = cudaMalloc((void**)&d_dist, sourceSize * sizeof(short));
    if (cudaStatus != cudaSuccess)
    {
        printf("cudaMalloc failed for histogram memory request!\n");
        cudaFree(d_source);
        return;
    }

    cudaStatus = cudaMemcpy(d_source, source, sourceSize * sizeof(short), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess)
    {
        printf("cudaMemcpy failed for sending model data on GPU!");
        cudaFree(d_source);
        cudaFree(d_dist);
    }

    cudaStatus = cudaMalloc((void**)&d_classes, classesCount * sizeof(int));
    if (cudaStatus != cudaSuccess)
    {
        printf("cudaMalloc failed for histogram memory request!\n");
        cudaFree(d_source);
        return;
    }

    cudaStatus = cudaMemcpy(d_classes, classes, classesCount * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess)
    {
        printf("cudaMemcpy failed for sending model data on GPU!");
        cudaFree(d_source);
        cudaFree(d_dist);
    }

    Timer t;
    cudaEvent_t start = t.go();

    transform_kernel<<<sourceSize/1024, 1024>>>(d_source, d_dist, d_classes,classesCount );

    cudaStatus = cudaDeviceSynchronize();

    printf("Matrix transform complete after : %f Ms \n", t.get(start));

    if (cudaStatus != cudaSuccess)
    {
        printf("cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    cudaStatus = cudaMemcpy(dist, d_dist, sourceSize * sizeof(short), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess)
    {
        printf("cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(d_source);
    cudaFree(d_dist);
}